-
Notifications
You must be signed in to change notification settings - Fork 220
New issue
Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.
By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.
Already on GitHub? Sign in to your account
Enable NCHW/NHWC and NCDHW/NDHWC layout in batch norm driver command #3234
base: develop
Are you sure you want to change the base?
Conversation
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
-L/--layout is quite generic. Which layout does it refer to? in/out/fil?
In convolution we do this: --fil_layout NCHW --in_layout NCHW --out_layout NCHW - which is nice bc its specific. I think BN should follow the same pattern.
Maybe there will be a time where in/fil layout in BN will be different, so we need to keep that in mind
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Some minor suggestions
driver/bn_driver.hpp
Outdated
} | ||
|
||
template <typename T> | ||
status_t AllocOnDevice(stream, context_t ctx, const size_t sz, std::vector<T>&) |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
That's suspicious:
Line 338 in 7e161f5
status_t AllocOnDevice(stream, context_t ctx, const size_t sz) |
I'm not sure that we want another copy.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
should I bring class GpumemTensor
from conv_driver.hpp
into another file driver_common.hpp
and let both conv_driver.hpp
and bn_driver.hpp
use it.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Yes, and somewhen, in the future, when you will see any raw or even managed gpu pointers in the driver\tests, you can recommend to use that class.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
For now I moved to driver.hpp
GpumemTensor
todriver/driver.hpp
miopenBNFwdTrainPerActivationRunHost
,miopenBNFwdTrainSpatialRunHost
,miopenBNFwdInferPerActivationRunHost
,miopenBNFwdInferSpatialRunHost
,miopenBNBwdPerActivationRunHost
andmiopenBNBwdSpatialRunHost
since they only handled NCHW layout. Started using layout agnostic batch norm host function fromtest/fusionHost.hpp
eg:
./bin/MIOpenDriver bnorm -n 64 -c 1024 -H 14 -W 14 -m 1 --forw 1 -b 0 -s 0 -r 1
// defaults to NCHW./bin/MIOpenDriver bnorm -n 64 -c 1024 -H 14 -W 14 -m 1 --forw 1 -b 0 -s 0 -r 1 -L NCHW
./bin/MIOpenDriver bnorm -n 64 -c 1024 -H 14 -W 14 -m 1 --forw 1 -b 0 -s 0 -r 1 -L NHWC
./bin/MIOpenDriver bnorm -n 64 -c 1024 -H 14 -W 14 -m 1 --forw 1 -b 0 -s 0 -r 1 --layout NHWC
bnorm, bnormfp16, bnormfp16fp32, bnormbfp16fp32,
3D
./bin/MIOpenDriver bnorm -n 8 -c 8 -H 12 -W 12 -D12 -m 1 --forw 0 -b 1 -s 1 -r 1 --layout NDHWC